#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cuda.h>
#include <device_functions.h>
#include <opencv2/opencv.hpp>
#include <iostream>
#include <math.h>
using namespace std;
using namespace cv;

//输入图像为BGR图，将其转化为gray图
__global__ void rgb2grayInCuda(uchar3 *dataIn, unsigned char *dataOut, int imgHeight, int imgWidth)
{
	int xIndex = threadIdx.x + blockIdx.x * blockDim.x;
	int yIndex = threadIdx.y + blockIdx.y * blockDim.y;

	if (xIndex < imgWidth && yIndex < imgHeight)
	{
		uchar3 rgb = dataIn[yIndex * imgWidth + xIndex];

		dataOut[yIndex * imgWidth + xIndex] = 0.299f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z;
	}
}

//灰度直方图统计
__global__ void imHistInCuda(unsigned char *dataIn, int *hist)
{
	int threadIndex = threadIdx.x + threadIdx.y * blockDim.x;
	int blockIndex = blockIdx.x + blockIdx.y * gridDim.x;
	int index = threadIndex + blockIndex * blockDim.x * blockDim.y;

	atomicAdd(&hist[dataIn[index]], 1);

}

int main()
{
	//传入图片
	Mat srcImg = imread("input.jpg");

	int imgHeight = srcImg.rows;
	int imgWidth = srcImg.cols;
	Mat grayImg(imgHeight, imgWidth, CV_8UC1, Scalar(0));//输出灰度图
	int hist[256];//灰度直方图统计数组
	memset(hist, 0, 256 * sizeof(int));

	//在GPU中开辟输入输出空间
	uchar3 *d_in;
	unsigned char *d_out;
	int *d_hist;

	cudaMalloc((void**)&d_in, imgHeight * imgWidth * sizeof(uchar3));
	cudaMalloc((void**)&d_out, imgHeight * imgWidth * sizeof(unsigned char));
	cudaMalloc((void**)&d_hist, 256 * sizeof(int));

	//将图像数据传入GPU中
	cudaMemcpy(d_in, srcImg.data, imgHeight * imgWidth * sizeof(uchar3), cudaMemcpyHostToDevice);
	cudaMemcpy(d_hist, hist, 256 * sizeof(int), cudaMemcpyHostToDevice);

	dim3 threadsPerBlock(32, 32);
	dim3 blocksPerGrid((imgWidth + threadsPerBlock.x - 1) / threadsPerBlock.x, (imgHeight + threadsPerBlock.y - 1) / threadsPerBlock.y);

	//灰度化
	rgb2grayInCuda << <blocksPerGrid, threadsPerBlock >> >(d_in, d_out, imgHeight, imgWidth);

	//灰度直方图统计
	imHistInCuda << <blocksPerGrid, threadsPerBlock >> >(d_out, d_hist);

	//将数据从GPU传回CPU
	cudaMemcpy(hist, d_hist, 256 * sizeof(int), cudaMemcpyDeviceToHost);
	cudaMemcpy(grayImg.data, d_out, imgHeight * imgWidth * sizeof(unsigned char), cudaMemcpyDeviceToHost);

    imwrite("output.jpg", grayImg);

	for(int i = 0; i < 256; i++) {
		std::cout << hist[i] << std::endl;
	}

	cudaFree(d_in);
	cudaFree(d_out);
	cudaFree(d_hist);

	return 0;
}